Unit_hipDeviceSynchronize_Functional used a host-memory read to assert the device was still busy before calling hipDeviceSynchronize(). That read raced with the in-flight D2H hipMemcpyAsync writing to the same pinned buffer — correctness depended entirely on the copy being slow enough to not finish first.

Motivation

Eliminate the data race in the pre-sync assertion of Unit_hipDeviceSynchronize_Functional. Reading shared host memory while a D2H async copy is in flight is undefined behavior; it was only masked by inflating the copy size via HIP_TEST_DEVICE_SYNCHRONIZE_FUNCTIONAL_COPY_MB, making the test fragile across fast GPUs and non-default HIP_LAUNCH_BLOCKING settings.

Technical Details

• projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc
  • Removed the racy REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1) pre-sync assertion and its associated workaround comment.
  • Replaced with a hipStreamQuery loop that checks for hipErrorNotReady — confirms at least one stream still has pending work without touching any shared host buffer:

// Before (racy): reads host buffer concurrent with D2H copy
REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);

// After: race-free stream status check
bool anyStreamBusy = false;
for (int i = 0; i < NUM_STREAMS; i++) {
  if (hipStreamQuery(stream[i]) == hipErrorNotReady) {
    anyStreamBusy = true;
    break;
  }
}
REQUIRE(anyStreamBusy);

• The post-sync correctness check (REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1)) is unchanged.

JIRA ID

ROCM-20519

Test Plan

Unit_hipDeviceSynchronize_Functional — verify the test passes and no longer requires oversized copies to mask timing races.

Test Result

No regressions observed in the modified test case.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

🔒 GitHub Advanced Security automatically protects Copilot coding agent pull requests. You can protect all pull requests by enabling Advanced Security for your repositories. Learn more about Advanced Security.